

#define CurrentID ((get_group_id(1)*get_num_groups(0)*get_num_groups(2)) + (get_group_id(2)*get_num_groups(0)) + get_group_id(0))
#define BlockNum 16

__kernel void f_ParticleDomain(__global float4* particle, __global ulong2* Hash, uint eNum, float gridGap){
	const size_t id = CurrentID;
	if(id<eNum){
		float3 p = particle[id].xyz;
		if(fabs(p.x) < 32767 && fabs(p.y) < 32767 && fabs(p.z) < 32767 ){
            
            p /= gridGap;
            p += 32767.0f;
			//ushort3 offset = (ushort3)(32767, 32767, 32767);
			ushort3 coord =  (ushort3)(p.x, p.y, p.z);
            
			//coord.x -= get_group_id(0)/3;
			//coord.y -= get_group_id(1)/3;
			//coord.z -= get_group_id(2)/3;
            
			Hash[id].s0 = upsample((uint)(coord.z), upsample(coord.y, coord.x));
            
			//printf("v4 = %v3f \n", p);
			//if(id==0){
				//ulong2 test = 0;
				//test.s0 = 0x1111<<32;
				//test.s0 = test.s0>>32;
				//printf("test = %#vhx ", test);
			//}
			
		}
	}

}
